const sampler_t sampler = CLK_ADDRESS_CLAMP_TO_EDGE |
CLK_FILTER_NEAREST;

__kernel void
Sobel_RGB(read_only image2d_t src, write_only image2d_t dst)
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    private float4 p00 = read_imagef(src, sampler, (int2)(x - 1, y - 1));
    private float4 p10 = read_imagef(src, sampler, (int2)(x, y - 1));
    private float4 p20 = read_imagef(src, sampler, (int2)(x + 1, y - 1));

    private float4 p01 = read_imagef(src, sampler, (int2)(x - 1, y));
    private float4 p21 = read_imagef(src, sampler, (int2)(x + 1, y));

    private float4 p02 = read_imagef(src, sampler, (int2)(x - 1, y + 1));
    private float4 p12 = read_imagef(src, sampler, (int2)(x, y + 1));
    private float4 p22 = read_imagef(src, sampler, (int2)(x + 1, y + 1));
    
    private float4 gx = -p00 + p20 +
        2.0f * (p21 - p01)
        -p02 + p22;
    private float4 gy = -p00 - p20 +
        2.0f * (p12 - p10) +
        p02 + p22;

    float4 g = native_sqrt(fma(gx,gx,fma(gy,gy,(float4)(0.0f))));
    // we could also approximate this as
    //float4 g = fabs(gx) + fabs(gy);
    // It creates basicly the same image but for exactness we will use the strict formula
    // I know that root(a^2+b^2) != a + b but It creates a similer image
    write_imagef(dst, (int2)(x, y), (float4)(g.xyz, 1.0f));
}